iT邦幫忙

2026 iThome 鐵人賽

DAY 23
0
Software Development

GPU 效能優化實戰:30 天從 Kernel 到 Profiling (重賽版)系列 第 23

Day 23:只移動一個 Node,為什麼要把整張圖重算一次?]] 重賽版[[

  • 分享至 

  • xImage
  •  

AI slob越來越進步 今天有點累 明天再mur mur


Day 22 的 Crossing evaluator 很老實。

如果 graph 有 E 條 edges,它讓每一個 GPU thread 負責一條 edge,再把其他 E 條 edges 全部掃過一次:

thread e
    -> crossing(e, 0)
    -> crossing(e, 1)
    -> crossing(e, 2)
    -> ...
    -> crossing(e, E - 1)
    -> output counts[e]

整體工作量是:

E * E

這個版本很適合當 correctness baseline,因為每一份 candidate layout 都從頭算,不需要相信任何舊狀態。

但 SA 一次通常只做這件事:

把一個 node v

從 old_position

移到 new_position

其他 nodes 沒有動,絕大多數 edges 也沒有動。

結果我們卻把整張圖的 edge pairs 又算了一次。

所以今天的優化不會先從 shared memory、warp shuffle 或 instruction unrolling 開始。我們先問一個更值錢的問題:

哪些 crossing 根本不可能改變?

Node 移動後,真正改變的只有 Incident Edges

假設 node v 連著三條 edges:

                  ● a
                 /
                / e0
               /
        b ●---● v---● c
            e1
                 \ e2
                  \
                   ● d

e0e1e2 都接到 node v,因此它們叫做 v 的 incident edges。

我們把這個集合寫成:

I(v) = all edges connected to node v

集合大小就是 node 的 degree:

d(v) = size(I(v))

v 被移動時,只有 I(v) 裡面的線段會改變形狀:

before move

                  ● a
                 /
                /
               /
              ● v

after move

                  ● a
              ___/
          ● v'

沒有接到 v 的 edge,兩個端點都沒動,所以線段也完全相同。

於是我們得到今天最重要的 invariant:

if edge e is not in I(v)
and edge f is not in I(v):

    crossing_old(e, f) == crossing_new(e, f)

這不是近似,也不是「通常不會改變」。只要兩條線段的四個端點都沒變,它們的 crossing 狀態就一定不變。


把所有 Edge Pairs 分成三組

移動 node v 後,任意一對 edges 只會落在下面三種情況。

第一種,兩條都是 static edges:

e not in I(v)
f not in I(v)

兩條線段都沒動
    -> crossing 不變
    -> 不用重算

第二種,一條是 incident edge,另一條是 static edge:

e in I(v)
f not in I(v)

其中一條線段改變
    -> crossing 可能出現
    -> crossing 也可能消失
    -> 必須比較 old 與 new

第三種,兩條都是 v 的 incident edges:

e in I(v)
f in I(v)

兩條 edges 共用 node v
    -> 不會被算成 proper crossing

第三種在 crossing count 中不會增加數量,但 geometry validator 仍然要檢查移動後是否出現共線重疊等非法情況。Crossing evaluator 少算工作,不代表合法性檢查可以一起刪掉。

把三組放在一起看:

                    partner edge f

                incident       static
              +------------+------------+
incident e    | shared v   | may change |
              +------------+------------+
static e      | may change | unchanged  |
              +------------+------------+

右下角的 static-static 區域通常最大,也是我們真正要刪掉的工作。


E * E 變成 d(v) * E

Day 22 的 full evaluator 不管移動哪個 node,都做:

work_full = E * E

增量更新只需要處理 incident edges 和其他 edges 的關係,主要工作量變成:

work_delta is proportional to d(v) * E

假設 graph 有 512 條 edges,而這次移動的 node degree 是 8:

full scale
    512 * 512
    = 262,144

affected-pair scale
    8 * 512
    = 4,096

這裡的 4,096 是用來理解數量級的 affected-pair scale,不是 kernel 實際執行的精確 instruction count。實作還需要比較 old/new crossing,並為 incident 與 static edges 各自產生正確的 per-edge count。

但方向已經很清楚:

d(v) << E

時,

d(v) * E << E * E

一般 graph 的 node degree 往往遠小於總 edge 數,因此演算法上少算,通常比把完整雙迴圈微調快一點更有價值。


Static Edge 可以直接從舊答案加上 Delta

先看一條沒有接到 moved node 的 static edge j

它和其他 static edges 的 crossing 狀態全部不變,所以不必從零開始計數。新的 count 可以從舊值修正:

new_count[j] = old_count[j]

接著只檢查它與每一條 incident edge i 的 old/new 狀態:

delta(i, j) = new_cross(i, j) - old_cross(i, j)

其中 crossing Boolean 可以當成 01

old = 0, new = 0 -> delta =  0
old = 0, new = 1 -> delta = +1
old = 1, new = 0 -> delta = -1
old = 1, new = 1 -> delta =  0

最後:

new_count[j] = old_count[j]
             + sum(
                   new_cross(i, j) - old_cross(i, j)
                   for i in I(v)
               )

這裡同時要算 old 與 new,不能只看 candidate layout。

如果一條 crossing 因為 node 移開而消失,只看 new layout 只會得到 0,卻不知道應該從舊 count 減掉 1


Incident Edge 為什麼乾脆重新計數?

再看一條接到 moved node 的 incident edge i

它本身的幾何位置已經改變,理論上它與任何 static edge 的 crossing 都可能不同。因此最清楚的做法是讓它從 0 開始,掃過所有 partner edges:

new_count[i] = 0

for every partner edge p:
    new_count[i] += new_cross(i, p)

這個 branch 仍然需要 E 等級的工作,但只有 d(v) 條 incident edges 會進來。

完整邏輯是:

for each output edge j:

    if j is incident to moved node:
        recount j against all partner edges

    else:
        begin with old_count[j]
        apply old/new delta from I(v)

我們仍然維持 Day 22 的 ownership:

一個 thread
    -> 負責一條 output edge
    -> 使用 local accumulator
    -> 最後寫一次 output

因此即使變成增量更新,也不需要讓許多 threads 對同一個 per-edge counter 做 atomicAdd


先建立 Node 到 Incident Edges 的索引

如果每個 proposal 都掃過全部 E 條 edges,詢問「這條 edge 有沒有接到 node v」,我們只是把成本藏到另一個迴圈裡。

Graph topology 在整個搜尋期間不會改變,所以 node-to-edge adjacency 應該只建立一次。

原專案使用 CSR 形式:

node_edge_offsets[V + 1]
incident_edges[2 * E]

例如:

edge 0 = (0, 1)
edge 1 = (0, 2)
edge 2 = (1, 3)
edge 3 = (2, 3)

每個 node 的 incident list 是:

node 0 -> [0, 1]
node 1 -> [0, 2]
node 2 -> [1, 3]
node 3 -> [2, 3]

壓平成兩個 arrays:

offsets  = [0, 2, 4, 6, 8]

incident = [
    0, 1,
    0, 2,
    1, 3,
    2, 3
]

要查 node 2,只需要:

begin = offsets[2]     = 4
end   = offsets[3]     = 6

incident[4:6] = [1, 3]

在 CUDA kernel 裡就是:

int begin = node_edge_offsets[moved_node];
int end   = node_edge_offsets[moved_node + 1];

for (int k = begin; k < end; ++k) {
    int edge_id = incident_edges[k];
    // process affected edge
}

為什麼 incident_edges 大小是 2 * E?因為一般 edge 有兩個端點,會分別出現在兩個 nodes 的 incident lists 裡。

這個索引在 CPU 建立一次,再上傳 GPU 一次。之後每個 proposal 都重用,不需要在 hot loop 裡重新配置或重新傳輸。

build once
    -> upload once
    -> reuse for millions of proposals

這個優化介紹的不是更快的乘法,而是資料結構如何決定 GPU 最後要做多少工作。


Candidate Move 不需要複製整份 Layout

假設 graph 有 N 個 nodes,我們同時評估 P 個 candidate moves。

最直接的資料表示可能是替每個 candidate 複製一份完整座標:

candidate_positions[P][N][2]

但一個 candidate 只移動一個 node。其餘 N - 1 組座標和 base layout 完全相同。

所以範例只保存三個整數:

move[p] = {
    node_id,
    new_x,
    new_y
}

讀取座標時才疊上 candidate:

int ax = (a == moved_node) ? new_x : base_x[a];
int ay = (a == moved_node) ? new_y : base_y[a];

如果有 128 個 nodes、32 個 proposals,使用 32-bit integer:

複製完整 layouts
    32 * 128 * 2 * 4 bytes
    = 32,768 bytes

只保存 moves
    32 * 3 * 4 bytes
    = 384 bytes

每批 proposals 額外需要的座標資料從 32 KB 降到 384 bytes;base layout 仍然只在 GPU 保存一份。graph 越大,差距越明顯。

這就是很典型的 GPU memory optimization:不是想辦法讓 PCIe 搬得更快,而是不要搬那些從來沒有改變的資料。


CUDA Kernel:一個 Proposal 乘上一排 Edges

教學範例同時評估 32 個 independent proposals。

Grid 可以想成一張二維表:

                    edge index
               0   1   2   3  ... E-1

proposal 0    [ ] [ ] [ ] [ ] ... [ ]
proposal 1    [ ] [ ] [ ] [ ] ... [ ]
proposal 2    [ ] [ ] [ ] [ ] ... [ ]
   ...
proposal P-1  [ ] [ ] [ ] [ ] ... [ ]

CUDA mapping 是:

int p = blockIdx.y;
int e = blockIdx.x * blockDim.x + threadIdx.x;

每個 cell 負責:

proposal p 的 edge e 新 crossing count

核心 kernel 可以縮成:

if (edge e is incident to moved_node) {
    count = 0;

    for (int f = 0; f < num_edges; ++f) {
        count += crossing_after_move(e, f);
    }
} else {
    count = base_count[e];

    for (int f : incident_edges[moved_node]) {
        count += crossing_after_move(e, f)
               - crossing_before_move(e, f);
    }
}

output[p * num_edges + e] = count;

E 很小時,單一 proposal 可能只有一兩個 blocks,根本填不滿 GPU。把多個 independent proposals 放進 blockIdx.y,可以提供更多 blocks,增加 throughput。

但這裡有一條不能偷換的語意:

32 個 independent proposals from one base layout

不等於

32 個 sequential SA steps

真正的 SA 如果接受 proposal 0,proposal 1 理論上應該從更新後的 layout 出發。教學 benchmark 中的 32 個 proposals 只是用來測試 evaluator throughput 與 correctness,沒有 commit moves,也沒有模擬一條完整 SA chain。


為什麼這個 Kernel 還是可能出現 Divergence?

同一個 warp 裡,大部分 edges 通常是 static,少數 edges 是 incident:

lane 0  -> static branch
lane 1  -> static branch
lane 2  -> incident branch
lane 3  -> static branch
...

incident branch 會掃 E 條 partners;static branch 只掃 d(v) 條 incident edges。兩邊工作量不一樣,因此 warp 可能需要等待較慢的 branch 完成。

這是增量演算法的代價之一:工作變少了,但規則不再像 full evaluator 那麼整齊。

進一步的實作可以把工作拆成不同 kernels:

kernel A
    更新 static candidate edges

kernel B
    重新計算 incident edges

這樣每個 kernel 裡的 threads 做比較一致的工作,但會多一次 kernel launch,也需要額外的 candidate list。是否划算,要看 graph 大小、degree 分布與 proposal batch size。

所以「避免 divergence」也不能脫離整個 pipeline 單獨判斷。


C1 還不夠:它少算 Pair,卻仍然啟動所有 Edge Threads

原專案把前面的 dense incremental scaffold 稱為 C1。

C1 已經把 pair tests 從 E * E 降到 d(v) * E 等級,但每個 proposal 仍會為所有 E 條 output edges 啟動 worker:

incident edge
    -> full recount

non-incident edge
    -> scan I(v) and apply delta

如果 graph 很大,下一個問題是:連 E 條 static edges 都有必要逐條看嗎?

答案仍然是否定的。

一條 incident segment 從舊位置移到新位置,只可能與空間上靠近它舊路徑或新路徑的 edges 改變 crossing 關係。這就是 C3 spatial candidate evaluator 的出發點。


C3:用 Spatial Grid 找可能改變的 Partner Edges

先把畫布切成格子,並記錄每一格有哪些 edges 經過:

+-------+-------+-------+-------+
|       | e7    |       |       |
+-------+-------+-------+-------+
| e2    | e2,e7 | e7    |       |
+-------+-------+-------+-------+
|       | e5    | e5    | e5    |
+-------+-------+-------+-------+

當 node v 移動時,對每一條 incident edge 取得:

old segment bounding box
new segment bounding box

再查詢兩者覆蓋的 grid cells:

old cells
    union
new cells
    -> candidate partner edges
    -> deduplicate
    -> exact crossing test

只有 candidate edges 需要計算 old/new delta。其他 non-incident、non-candidate edges 可以直接複製舊 count:

new_count[j] = old_count[j]

Spatial grid 只負責縮小 candidate set,最後的答案仍然由 exact integer crossing predicate 決定。它不是用格子近似 crossing。


反直覺的地方:Old Region 和 New Region 都要查

很多人第一次寫 spatial delta,會只查新線段經過的 cells:

new bounding box
    -> 找現在可能 crossing 的 edges

這只能抓到新產生的 crossing,卻可能漏掉已經消失的 crossing。

例如:

before move

incident edge i  X  static edge j
    old_cross(i, j) = 1

after move

incident edge i     static edge j
    new_cross(i, j) = 0

如果 j 只出現在 old segment 的區域,查詢 new region 時根本看不到它。counts[j] 就不會減掉 1,錯誤會留在 state 裡。

因此 candidate set 必須來自:

candidates = edges_in(old_region)
             union
             edges_in(new_region)

這是增量演算法很常見的陷阱:新狀態告訴你增加了什麼,舊狀態才告訴你失去了什麼。


Candidate List 滿了,不能偷偷截斷

Spatial grid 通常使用固定容量,因為我們不想在每個 proposal 裡動態配置記憶體。

但固定容量一定會遇到 overflow:

某個 cell 裡的 edges 太多

或

union 後的 candidate edges 太多

最危險的做法是只保留前 capacity 條 edges:

candidate_count > capacity
    -> silently truncate

這會讓 kernel 很快,答案卻不再 exact。

原專案的正確策略是:

overflow detected
    -> 記錄 typed fallback reason
    -> 放棄這次不完整的 candidate list
    -> 回到 dense C1 或 full-square path
    -> 只發布 exact result

fallback 偶爾比較慢沒有關係。對 optimizer 來說,罕見的慢答案仍然比無聲的錯答案好;錯誤 crossing counts 一旦被 accepted move 寫回 current state,後面的 SA 會在錯誤狀態上繼續跑,而且可能很久才被發現。


Generation Stamp:連清空 Flag Array 都可以省

C3 需要標記哪些 edges 是本次 candidates。最直接的方式是:

is_candidate[E] = all zeros

for edge in candidates:
    is_candidate[edge] = 1

但這代表每個 proposal 都要先清空整個 E 長度的 array。

另一種方法是使用 generation stamp:

current_generation = 137

candidate_generation[edge] = 137

判斷時比較:

candidate_generation[edge] == current_generation

下一次 proposal 只要把 generation 加一,不必先把整個 array 歸零:

proposal 137 -> generation 137
proposal 138 -> generation 138
proposal 139 -> generation 139

這個技巧省掉的是固定的 memory write。當我們已經努力把 crossing tests 減少後,這類每回合都掃完整 array 的初始化成本就會開始變得顯眼。

當 generation integer 即將 wrap 時,仍需要安全地 reset;這也是實作不能省略的邊界條件。


效能結果:少算很多,為什麼只快 1.96 倍?

教學範例使用 RTX 4060 Laptop GPU,固定 seed = 42,一次 kernel 評估 32 個 independent proposals。

結果如下:

Graph Full median Dense delta median Full / Delta
48 nodes / 96 edges 0.02473 ms 0.02243 ms 1.10x
128 nodes / 512 edges 0.25788 ms 0.13190 ms 1.96x

這裡的 delta 是本篇前半段的 dense incremental C1 教學 kernel,沒有包含 C3 spatial grid。

測量範圍是:

CUDA events
5 次 warm-up
9 batches
每個 batch 連續執行 100 次
最後取 median

包含的是 GPU stream 上的 kernel 執行區間。不包含:

CuPy JIT compile
memory allocation
host-device transfer
Python CPU oracle
完整 V1~V4 validator
SA acceptance 與 commit

512-edge case 的 crossing work 在演算法上減少很多,實測卻只有大約 1.96x。這其實很合理:

第一,C1 仍然為每條 output edge 啟動 thread

第二,static branch 要同時計算 old 與 new crossing

第三,CSR incident list 帶來額外的間接讀取

第四,incident 與 static branch 工作量不同,可能產生 divergence

第五,output 仍然是 proposals * edges 個 counts

第六,小 kernel 仍受固定 launch 與排程成本影響

96-edge case 只快 1.10x,也不是優化失敗。工作太小時,省下來的幾何測試還不足以壓過固定成本。

這是本篇最值得保留的 GPU 經驗:

演算法複雜度下降

不代表

小資料一定立刻得到同等比例的 speedup

資料量、parallelism、memory access 與固定成本,最後都會反映在 wall time 裡。


Correctness:Delta 最怕的不是單次算錯,而是 State Drift

Full recompute 每次從座標重新建立 counts。Delta evaluator 則依賴:

current coordinates
current per-edge counts
candidate move

如果某次更新漏掉一條 edge,錯誤會被寫進 current counts。下一次 proposal 又以錯誤 counts 當 base,偏差可能一路累積。

因此測試不能只看最後的 K,必須逐 edge 比較:

delta_counts(candidate)
    ==
full_recompute(candidate)

至少要覆蓋:

移動後新增一個 crossing
移動後移除一個 crossing
crossing pair 的兩條 edges 都正確更新
完全不受影響的 edge count 保持不變
degree 0 node
degree 1 node
high-degree node
重複接受多次 moves 後仍然沒有 drift
spatial old region 能抓到消失的 crossing
candidate overflow 正確 fallback

教學範例會讓 GPU full counts、GPU delta counts 與獨立 CPU oracle 逐 edge 相等,才接受結果。


真正的 SA 還多一個 Transactional Commit 問題

範例中的 proposals 都從同一份 immutable base layout 出發,因此很容易驗證:不需要把任何 candidate 寫回 current state。

真正的 SA 則必須分成兩個階段:

evaluate
    把 candidate coordinates 與 counts 寫到 scratch
    計算 candidate objective

accept
    才把 coordinates、counts、objective 一起 commit

拒絕 proposal 時:

current coordinates unchanged
current counts unchanged
current objective unchanged

接受 proposal 時,這三份狀態必須一起前進:

new coordinates
new per-edge counts
new K, n_K, Phi, C

不能讓其他 kernel 看見「新座標配舊 counts」,也不能先發布新的 objective、稍後才補 crossing state。在單一 block 內可以用 cooperative phases 與 synchronization;跨 kernels 則要靠同一 stream 的 ordering 保證提交順序。

這部分不會出現在單獨的 crossing microbenchmark 裡,卻是把演算法放進完整 GPU solver 時最容易踩到的工程問題之一。


Per-edge Counts 算完,不代表 Objective 免費完成

本篇 benchmark 只量 crossing-count kernel。

得到新的 counts[e] 之後,solver 還要算:

K
n_K
Phi
C

最直接的方法是對完整 E 長度的 scratch counts 做 reduction。這已經比 E * E 的 full crossing evaluator 便宜很多,但仍然是一筆成本。

更深的版本可以維護 crossing-count histogram,只更新真正改變的 buckets,再從 non-empty buckets 找新的 K。代價是 state 與 commit 邏輯會更複雜。

因此這篇的 1.96x 只能回答:

dense full crossing-count kernel

和

dense incremental crossing-count kernel

哪一個比較快?

它不能回答完整 SA solver 最後快了幾倍。要得到那個答案,必須把 validation、objective reduction、proposal generation、acceptance、commit 與同步全部放進相同的 measurement scope。


今天真正學到的 GPU 優化

Day 23 做的不是把 Day 22 kernel 換成比較花俏的 CUDA 語法,而是重新定義需要執行的工作:

  1. 只有 moved node 的 incident edges 會改變幾何。
  2. Static edge 從舊 count 加上 old/new delta,不必完整重算。
  3. Incident edge 仍做 exact recount,維持單一 writer ownership。
  4. Node-to-edge CSR 建立一次並常駐 GPU,避免每回合找 adjacency。
  5. Candidate 只保存 (node, new_x, new_y),不複製完整 layout。
  6. 多個 independent proposals 提供 throughput,但不能假裝成 sequential SA steps。
  7. Spatial grid 必須查 old/new region 的 union,才能同時處理新增與消失的 crossings。
  8. 固定容量 overflow 必須 fallback,不能為了速度截斷 exact candidate set。
  9. Kernel microbenchmark 與完整 solver speedup 是不同問題。

Day 22 的問題是:

怎麼把所有 crossing counts 算對?

Day 23 的問題則是:

既然只移動一個 node,
怎麼保留不變的答案,只更新真的可能改變的部分?

下一篇會進入 Layout。Crossing evaluator 已經告訴我們一份 layout 有多好,接下來要問的是:第一份 node positions 從哪裡來?為什麼隨機亂放雖然簡單,卻可能讓後面的 SA 一開始就背著大量 crossings?

程式對照


上一篇
Day 22:先把 Crossing 算對,再談怎麼少算 ))重賽版((
下一篇
Day 24:Crossing Evaluator 會打分數,但第一張圖從哪裡來?}}重賽版{{
系列文
GPU 效能優化實戰:30 天從 Kernel 到 Profiling (重賽版)26
圖片
  熱門推薦
圖片
{{ item.channelVendor }} | {{ item.webinarstarted }} |
{{ formatDate(item.duration) }}
直播中

尚未有邦友留言

立即登入留言